#include "CCL_gpu.h"

#include <assert.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <helper_cuda.h>
#include <device_functions.h> //_synchronizedThreads()
#include <stdlib.h>
#include <assert.h>
#include <unistd.h>
#define STANDAR_DIMBLOCK 32

/**
 * PRE: width > 0 && height > 0 && pixelsPerThread > 0 && dims != NULL && deviceProp != NULL
 * @param dims array of integers with room for six of them
 * @param deviceProp structure cudaDeviceProp
 * @param width integer that represents the width of the image
 * @param height integer that represents the height of the image
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 */
void calculateDimSmem(int* dims, cudaDeviceProp deviceProp, int width, int height, int pixelsPerThread){

	//THREADS POR BLOQUE
	dims[0] = sqrt(deviceProp.maxThreadsPerBlock)/pixelsPerThread;

	dims[1] = deviceProp.maxThreadsPerBlock/dims[0];



	//MEMORIA SHARED
	dims[4] = pixelsPerThread * dims[0];
	dims[5] = dims[1];
	//BLOQUES EN LA GRILLA

	dims[2] = (width + (dims[4]-1))/dims[4];
	dims[3] = (height + (dims[5]-1))/dims[5];
}
/**
 *
 * @param newLabel
 * @param neighbor
 */
inline __device__ void compareNeighbor(int* newLabel, int neighbor){
	 if( neighbor < *newLabel && neighbor!=0){
		 *newLabel = neighbor;
	 }
}

/**
 * Search int the matrix buf, the root of the label x and returns root's position value
 * PRE:buf!=NULL && x>=0
 * @param buf integer's array
 * @param x position in buf
 * POS: result = buf[result]
 * invariant: nextX>=0 && buf[nextX]>=0
 * @return the position of the root
 */
inline __device__ int findRoot(int* buf, int x) {
	int next=x;
	int result=-1;
    while (result != next){
    	 result = next;
    	 next  = buf[next]-1;
    }
    return result;
}
/*
 *
 *
 *
 */
/**
 * Replace one matrix's label with his root
 * PRE: x>=0 && y>=0 && bufOut!=NULL && bufIn!=NULL
 * @param x width's position in bufIn
 * @param y height's position in bufIn
 * @param bufOut integer's array
 * @param bufIn integer's array
 * POS:label <= findRoot(bufIn,label)
 * invariant: newLabel>=0
 */
inline __device__ void updatePixel(int x, int y,int* bufOut,int* bufIn){
	int index = x+y;
	int label = bufIn[index]-1;
	if(label != index && label >= 0){
		int newLabel = findRoot(bufIn,label);
		assert(newLabel>=0);
			bufOut[index] = newLabel+1;
	}
}
/*
 *
 */
/**
 * Replace all matrix's elements with them root
 * PRE: width>0 && y!=0 && dLabelsOut!=NULL && dLabelsIn!=NULL
 * @param bufOut integer's array
 * @param bufIn integer's array
 * @param width integer that represents the width of the image
 * @param height integer that represents the height of the image
 * @param pixelsPerThread integer that represents how match pixels a thread has to work with
 * @param dimX length's dimension X in the block
 * @param dimY length's dimension Y in the block
 */
__global__ void update(int* buf, int width, int height, int pixelsPerThread, int dimX, int dimY){
	//ME DICE EN QUE BLOQUE ESTA
	int Y = blockIdx.y*dimY+threadIdx.y;
	int idxX =dimX/pixelsPerThread;

	for(int X=threadIdx.x;X<dimX;X=X+idxX){
		int tmp = blockIdx.x*dimX+X;
		if(tmp<width && Y<height){
			updatePixel(tmp, Y*width, buf, buf);
		}
	}
}

__device__ void setDimblock(int* dimblock,int width,int height, int dimX, int dimY){
	int tmp=0;
	if(blockIdx.x==gridDim.x){
		tmp=width%dimX;
		if(tmp!=0){
			dimblock[0]=tmp;
		}
		else{
			dimblock[0]=dimX;
		}
	}
	else{
		dimblock[0]=dimX;
	}
	if(blockIdx.y==gridDim.y){
		tmp=height%dimY;
		if(tmp!=0){
			dimblock[1]=tmp;
		}
		else{
			dimblock[1]=dimY;
		}

	}
	else{
			dimblock[1]=dimY;
		}

}

/**
 * Divide the matrix in smaller ones and calculate the CCL for each one
 * PRE:d_data!=NULL && width>=0 && height>=0
 * @param d_data
 * @param width
 * @param height
 * @param pixelsPerThread
 * @param dimX
 * @param dimY
 *
 */
__global__ void LocalSolution(uchar4* d_img,int* d_data,int width,int height, int pixelsPerThread, int dimX, int dimY,int umbral){
	//__shared__ bool change[1];
	//extern __shared__ int sLabels[];

	//ME DICE EN QUE BLOQUE ESTA
	int Y = blockIdx.y*dimY+threadIdx.y;
	int oldLabel;
	int idxX =dimX/pixelsPerThread;
	int dimblock[2]={0};
	setDimblock(dimblock,width,height, dimX,dimY);
	if(Y<height){
		for(int X=threadIdx.x;X<dimX;X=X+idxX){
			int tmp = blockIdx.x*dimX+X;
			if(tmp<width){
				uchar4 rgb = d_img[Y*width+tmp];
				if((float)((int)rgb.x + (int)rgb.y + (int)rgb.z)/3 >= umbral){
				d_data[Y*width + tmp] = Y*width+tmp+1;
					oldLabel =Y*width+tmp+1;
				}
				else{
					d_data[Y*width + tmp] = 0;
					oldLabel = 0;
				}
				//sLabels[threadIdx.y*dimX + X]=oldLabel;
			}
		}

	/*if(threadIdx.y<dimY){
		change[0]=1;
		__syncthreads();
		while(change[0]){
			change[0]=0;
			for(int X=threadIdx.x;X<dimX;X=X+idxX){
				oldLabel = sLabels[threadIdx.y*dimX + X];
				if(oldLabel!=0){
					int newLabel[1]={oldLabel};
					int xMin = max(X-1,0);
					int xMax = min(X+1,dimblock[0]-1);
					int yMin = max((int)threadIdx.y-1,0)*dimX;
					int yMax = min((int)threadIdx.y+1,dimblock[1]-1)*dimX;

					compareNeighbor(newLabel,sLabels[xMin + yMin]);
					compareNeighbor(newLabel,sLabels[xMin + threadIdx.y*dimX]);
					compareNeighbor(newLabel,sLabels[xMin + yMax]);
					compareNeighbor(newLabel,sLabels[X + yMin]);
					compareNeighbor(newLabel,sLabels[X + yMax]);

					compareNeighbor(newLabel,sLabels[xMax + yMin]);
					compareNeighbor(newLabel,sLabels[xMax + threadIdx.y*dimX]);
					compareNeighbor(newLabel,sLabels[xMax + yMax]);
					if(oldLabel > *newLabel) {
					/*	int root = findRoot(sLabels,threadIdx.y*dimX + X);
						assert(root>=0);
						sLabels[root]=min(sLabels[root],*newLabel);
						sLabels[threadIdx.y*dimX + X]=*newLabel;
						change[0]=1;
					}
				}
			}
			__syncthreads();
		}*/
/*		for(int X=threadIdx.x;X<dimX;X=X+idxX){
			int tmp = blockIdx.x*dimX+X;
			if(tmp<width){
				d_data[Y*width + tmp]=sLabels[threadIdx.y*dimX + X];
			}
		}*/
	}
}

__global__ void mergeAll(int* d_data,int width,int height,bool* d_changed){

	int dimX = blockDim.x;
	int dimY = blockDim.y;
	int Y = blockIdx.y*dimY+threadIdx.y;
	int X = blockIdx.x*dimX+threadIdx.x;


	if(X<width && Y<height){
		int oldLabel = d_data[Y*width + X];
		if(oldLabel!=0){
			int newLabel[1]={oldLabel};
			int xMin = max(X-1,0);
			int xMax = min(X+1,width-1);
			int yMin = max(Y-1,0)*width;
			int yMax = min(Y+1,height-1)*width;

			compareNeighbor(newLabel,d_data[xMin + yMin]);
			compareNeighbor(newLabel,d_data[xMin + Y*width]);
			compareNeighbor(newLabel,d_data[xMin + yMax]);

			compareNeighbor(newLabel,d_data[X + yMin]);
			compareNeighbor(newLabel,d_data[X + yMax]);

			compareNeighbor(newLabel,d_data[xMax + yMin]);
			compareNeighbor(newLabel,d_data[xMax + Y*width]);
			compareNeighbor(newLabel,d_data[xMax + yMax]);
			__syncthreads();
			if(oldLabel > *newLabel) {
				int root = findRoot(d_data,X + Y*width);
				assert(root>=0);
				d_data[root]=min(d_data[root],*newLabel);
				*d_changed = 1;
			}
		}
	}
}

/**
 * PRE:h_img!=NULL width!=0 height!=0
 * @param h_img
 * @param width
 * @param height
 * @param dims
 * @param pixelsPerThread
 * @param umbral
 * @return matrix CCL
 */
int* CCL_gpu(uchar4* h_img, int width, int height, cudaDeviceProp deviceProp, int pixelsPerThread, int umbral){
	assert(h_img!=NULL && width>0 && height>0);

	uchar4* d_img = NULL;
	int* h_data = NULL;
	int*d_data = NULL;
	bool* h_changed = NULL;
	bool* d_changed = NULL;
	int sMem = 0;
	int dims[6] = {0};
	calculateDimSmem(dims, deviceProp, width, height, pixelsPerThread);

	checkCudaErrors(cudaMalloc((void**) &(d_data), width*height*sizeof(int)));
	checkCudaErrors(cudaMalloc((void**) &(d_changed), sizeof(bool)));
	checkCudaErrors(cudaMalloc((void**) &(d_img), width*height*sizeof(uchar4)));

	cudaHostAlloc((void**)&h_data, width*height*sizeof(int),cudaHostAllocDefault);
	//checkCudaErrors(cudaMemcpy(d_data, h_data, width*height*sizeof(int), cudaMemcpyHostToDevice));

	cudaHostAlloc((void**)&h_changed,sizeof(bool),cudaHostAllocDefault);

	checkCudaErrors(cudaMemcpy(d_img, h_img, width*height*sizeof(uchar4), cudaMemcpyHostToDevice));
	//CREA GRILLA Y BLOQUES
	dim3 dimGrid((width+STANDAR_DIMBLOCK-1) /STANDAR_DIMBLOCK,(height+STANDAR_DIMBLOCK-1)/STANDAR_DIMBLOCK,1);
	dim3 dimGrid1(dims[2],dims[3],1);
	dim3 threads(STANDAR_DIMBLOCK,STANDAR_DIMBLOCK,1);
	dim3 threads1(dims[0],dims[1],1);
	sMem = dims[4]*dims[5]*sizeof(int);



		//printf("LocalSolution...\n");
		LocalSolution<<<dimGrid1,threads1,sMem>>>(d_img,d_data,width,height,pixelsPerThread, dims[4], dims[5],umbral);
		int i=0;
		do{
			i++;
			cudaMemset((void*) d_changed,0,sizeof(bool));
		//	printf("mergeAll...\n");
			mergeAll<<<dimGrid,threads>>>(d_data,width,height,d_changed);

			checkCudaErrors(cudaMemcpy(h_changed, d_changed, sizeof(bool), cudaMemcpyDeviceToHost));
		//	printf("update...\n");
			update<<<dimGrid1,threads1>>>(d_data,width,height, pixelsPerThread, dims[4], dims[5]);

		}while(*h_changed);
		printf("iteraciones while : %d\n",i);
/*	*/
	//COPIADO DE MEMORIA Y LIBERACION DE RECURSOS
	checkCudaErrors(cudaMemcpy(h_data, d_data, width*height*sizeof(int), cudaMemcpyDeviceToHost));
	checkCudaErrors(cudaFree(d_img));
	checkCudaErrors(cudaFree(d_data));
	checkCudaErrors(cudaFree(d_changed));
	//checkCudaErrors(cudaStreamDestroy(stream1));
	cudaFreeHost(h_changed);




	return h_data;
}
